<!-- HTML header for doxygen 1.8.3.1-->
<!DOCTYPE html PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml">
<head>
<meta http-equiv="Content-Type" content="text/xhtml;charset=UTF-8"/>
<meta http-equiv="X-UA-Compatible" content="IE=9"/>
<meta name="generator" content="Doxygen 1.8.4"/>
<title>CUB: cub Namespace Reference</title>
<link href="tabs.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="jquery.js"></script>
<script type="text/javascript" src="dynsections.js"></script>
<link href="search/search.css" rel="stylesheet" type="text/css"/>
<script type="text/javascript" src="search/search.js"></script>
<script type="text/javascript">
  $(document).ready(function() { searchBox.OnSelectItem(0); });
</script>
<link href="doxygen.css" rel="stylesheet" type="text/css" />
<link href="extra_stylesheet.css" rel="stylesheet" type="text/css"/>
<link rel="shortcut icon" href="favicon.ico" type="image/x-icon" />
<script type="text/javascript">
  var _gaq = _gaq || [];
  _gaq.push(['_setAccount', 'UA-38890655-1']);
  _gaq.push(['_trackPageview']);
  (function() {
    var ga = document.createElement('script'); ga.type = 'text/javascript'; ga.async = true;
    ga.src = ('https:' == document.location.protocol ? 'https://ssl' : 'http://www') + '.google-analytics.com/ga.js';
    var s = document.getElementsByTagName('script')[0]; s.parentNode.insertBefore(ga, s);
  })();
</script>
</head>
<body>
<div id="top"><!-- do not remove this div, it is closed by doxygen! -->
<div id="titlearea">
<table cellspacing="0" cellpadding="0">
 <tbody>
 <tr style="height: 56px;">
  <td style="padding-left: 0.5em;">
   <div id="projectname">CUB
   </div>
  </td>
 </tr>
 </tbody>
</table>
</div>
<!-- end header part -->
<!-- Generated by Doxygen 1.8.4 -->
<script type="text/javascript">
var searchBox = new SearchBox("searchBox", "search",false,'Search');
</script>
  <div id="navrow1" class="tabs">
    <ul class="tablist">
      <li><a href="index.html"><span>Main&#160;Page</span></a></li>
      <li><a href="modules.html"><span>Modules</span></a></li>
      <li><a href="annotated.html"><span>Classes</span></a></li>
      <li>
        <div id="MSearchBox" class="MSearchBoxInactive">
        <span class="left">
          <img id="MSearchSelect" src="search/mag_sel.png"
               onmouseover="return searchBox.OnSearchSelectShow()"
               onmouseout="return searchBox.OnSearchSelectHide()"
               alt=""/>
          <input type="text" id="MSearchField" value="Search" accesskey="S"
               onfocus="searchBox.OnSearchFieldFocus(true)" 
               onblur="searchBox.OnSearchFieldFocus(false)" 
               onkeyup="searchBox.OnSearchFieldChange(event)"/>
          </span><span class="right">
            <a id="MSearchClose" href="javascript:searchBox.CloseResultsWindow()"><img id="MSearchCloseImg" border="0" src="search/close.png" alt=""/></a>
          </span>
        </div>
      </li>
    </ul>
  </div>
  <div id="navrow2" class="tabs2">
    <ul class="tablist">
      <li><a href="namespaces.html"><span>Namespace&#160;List</span></a></li>
      <li><a href="namespacemembers.html"><span>Namespace&#160;Members</span></a></li>
    </ul>
  </div>
<!-- window showing the filter options -->
<div id="MSearchSelectWindow"
     onmouseover="return searchBox.OnSearchSelectShow()"
     onmouseout="return searchBox.OnSearchSelectHide()"
     onkeydown="return searchBox.OnSearchSelectKey(event)">
<a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(0)"><span class="SelectionMark">&#160;</span>All</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(1)"><span class="SelectionMark">&#160;</span>Classes</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(2)"><span class="SelectionMark">&#160;</span>Namespaces</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(3)"><span class="SelectionMark">&#160;</span>Files</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(4)"><span class="SelectionMark">&#160;</span>Functions</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(5)"><span class="SelectionMark">&#160;</span>Variables</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(6)"><span class="SelectionMark">&#160;</span>Typedefs</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(7)"><span class="SelectionMark">&#160;</span>Enumerations</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(8)"><span class="SelectionMark">&#160;</span>Enumerator</a><a class="SelectItem" href="javascript:void(0)" onclick="searchBox.OnSelectItem(9)"><span class="SelectionMark">&#160;</span>Groups</a></div>

<!-- iframe showing the search results (closed by default) -->
<div id="MSearchResultsWindow">
<iframe src="javascript:void(0)" frameborder="0" 
        name="MSearchResults" id="MSearchResults">
</iframe>
</div>

</div><!-- top -->
<div class="header">
  <div class="summary">
<a href="#nested-classes">Classes</a> &#124;
<a href="#enum-members">Enumerations</a> &#124;
<a href="#func-members">Functions</a>  </div>
  <div class="headertitle">
<div class="title">cub Namespace Reference</div>  </div>
</div><!--header-->
<div class="contents">

<p>Optional outer namespace(s)  
<a href="#details">More...</a></p>
<table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="nested-classes"></a>
Classes</h2></td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_arch_props.html">ArchProps</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Structure for statically reporting CUDA device properties, parameterized by SM architecture.  <a href="structcub_1_1_arch_props.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_array_wrapper.html">ArrayWrapper</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A wrapper for passing simple static arrays as kernel parameters.  <a href="structcub_1_1_array_wrapper.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_base_traits.html">BaseTraits</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Basic type traits.  <a href="structcub_1_1_base_traits.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_binary_op_has_idx_param.html">BinaryOpHasIdxParam</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Determine whether or not BinaryOp's functor is of the form <code>bool operator()(const T&amp; a, const T&amp;b)</code> or <code>bool operator()(const T&amp; a, const T&amp;b, unsigned int idx)</code>  <a href="structcub_1_1_binary_op_has_idx_param.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_discontinuity.html">BlockDiscontinuity</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_discontinuity.html" title="The BlockDiscontinuity class provides collective methods for flagging discontinuities within an order...">BlockDiscontinuity</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for flagging discontinuities within an ordered set of items partitioned across a CUDA thread block. </p>
<div class="image">
<img src="discont_logo.png" alt="discont_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_discontinuity.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_exchange.html">BlockExchange</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">BlockExchange</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for rearranging data partitioned across a CUDA thread block. </p>
<div class="image">
<img src="transpose_logo.png" alt="transpose_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_exchange.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_histogram.html">BlockHistogram</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_histogram.html" title="The BlockHistogram class provides collective methods for constructing block-wide histograms from data...">BlockHistogram</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for constructing block-wide histograms from data samples partitioned across a CUDA thread block. </p>
<div class="image">
<img src="histogram_logo.png" alt="histogram_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_histogram.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_load.html">BlockLoad</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">BlockLoad</a> class provides <a href="index.html#sec0"><em>collective</em></a> data movement methods for loading a linear segment of items from memory into a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> across a CUDA thread block. </p>
<div class="image">
<img src="block_load_logo.png" alt="block_load_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_load.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_radix_sort.html">BlockRadixSort</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_radix_sort.html" title="The cub::BlockRadixSort class provides collective methods for sorting items partitioned across a CUDA...">cub::BlockRadixSort</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for sorting items partitioned across a CUDA thread block using a radix sorting method. </p>
<div class="image">
<img src="sorting_logo.png" alt="sorting_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_radix_sort.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_reduce.html">BlockReduce</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_reduce.html" title="The BlockReduce class provides collective methods for computing a parallel reduction of items partiti...">BlockReduce</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel reduction of items partitioned across a CUDA thread block. </p>
<div class="image">
<img src="reduce_logo.png" alt="reduce_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_reduce.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_scan.html">BlockScan</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">BlockScan</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel prefix sum/scan of items partitioned across a CUDA thread block. </p>
<div class="image">
<img src="block_scan_logo.png" alt="block_scan_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_scan.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_block_store.html">BlockStore</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">BlockStore</a> class provides <a href="index.html#sec0"><em>collective</em></a> data movement methods for writing a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of items partitioned across a CUDA thread block to a linear segment of memory. </p>
<div class="image">
<img src="block_store_logo.png" alt="block_store_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_block_store.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_caching_device_allocator.html">CachingDeviceAllocator</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple caching allocator for device memory allocations.  <a href="structcub_1_1_caching_device_allocator.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_cast.html">Cast</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default cast functor.  <a href="structcub_1_1_cast.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_constant_iterator_r_a.html">ConstantIteratorRA</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple random-access iterator pointing to a range of constant values.  <a href="classcub_1_1_constant_iterator_r_a.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_device.html">Device</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Properties of a given CUDA device and the corresponding PTX bundle.  <a href="classcub_1_1_device.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_device_histogram.html">DeviceHistogram</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="structcub_1_1_device_histogram.html" title="DeviceHistogram provides device-wide parallel operations for constructing histogram(s) from samples d...">DeviceHistogram</a> provides device-wide parallel operations for constructing histogram(s) from samples data residing within global memory. </p>
<div class="image">
<img src="histogram_logo.png" alt="histogram_logo.png"/>
<div class="caption">
.</div></div>
  <a href="structcub_1_1_device_histogram.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_device_radix_sort.html">DeviceRadixSort</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="structcub_1_1_device_radix_sort.html" title="DeviceRadixSort provides operations for computing a device-wide, parallel radix sort across data item...">DeviceRadixSort</a> provides operations for computing a device-wide, parallel radix sort across data items residing within global memory. </p>
<div class="image">
<img src="sorting_logo.png" alt="sorting_logo.png"/>
<div class="caption">
.</div></div>
  <a href="structcub_1_1_device_radix_sort.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_device_reduce.html">DeviceReduce</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="structcub_1_1_device_reduce.html" title="DeviceReduce provides operations for computing a device-wide, parallel reduction across data items re...">DeviceReduce</a> provides operations for computing a device-wide, parallel reduction across data items residing within global memory. </p>
<div class="image">
<img src="reduce_logo.png" alt="reduce_logo.png"/>
<div class="caption">
.</div></div>
  <a href="structcub_1_1_device_reduce.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_device_scan.html">DeviceScan</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="structcub_1_1_device_scan.html" title="DeviceScan provides operations for computing a device-wide, parallel prefix scan across data items re...">DeviceScan</a> provides operations for computing a device-wide, parallel prefix scan across data items residing within global memory. </p>
<div class="image">
<img src="device_scan.png" alt="device_scan.png"/>
<div class="caption">
.</div></div>
  <a href="structcub_1_1_device_scan.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_double_buffer.html">DoubleBuffer</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Double-buffer storage wrapper for multi-pass stream transformations that require more than one storage array for streaming intermediate results back and forth.  <a href="structcub_1_1_double_buffer.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_enable_if.html">EnableIf</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Simple enable-if (similar to Boost)  <a href="structcub_1_1_enable_if.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_equality.html">Equality</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default equality functor.  <a href="structcub_1_1_equality.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_equals.html">Equals</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Type equality test.  <a href="structcub_1_1_equals.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_if.html">If</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Type selection (<code>IF ? ThenType : ElseType</code>)  <a href="structcub_1_1_if.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_inequality.html">Inequality</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default inequality functor.  <a href="structcub_1_1_inequality.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_int2_type.html">Int2Type</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Allows for the treatment of an integral constant as a type at compile-time (e.g., to achieve static call dispatch based on constant integral values)  <a href="structcub_1_1_int2_type.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_is_pointer.html">IsPointer</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Pointer vs. iterator.  <a href="structcub_1_1_is_pointer.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_is_volatile.html">IsVolatile</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Volatile modifier test.  <a href="structcub_1_1_is_volatile.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_log2.html">Log2</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Statically determine log2(N), rounded up.  <a href="structcub_1_1_log2.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_max.html">Max</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default max functor.  <a href="structcub_1_1_max.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_min.html">Min</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default min functor.  <a href="structcub_1_1_min.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_null_type.html">NullType</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple "NULL" marker type.  <a href="structcub_1_1_null_type.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_numeric_traits.html">NumericTraits</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Numeric type traits.  <a href="structcub_1_1_numeric_traits.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_power_of_two.html">PowerOfTwo</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Statically determine if N is a power-of-two.  <a href="structcub_1_1_power_of_two.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_ptx_arch_props.html">PtxArchProps</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The architectural properties for the PTX version targeted by the active compiler pass.  <a href="structcub_1_1_ptx_arch_props.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_remove_qualifiers.html">RemoveQualifiers</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Removes <code>const</code> and <code>volatile</code> qualifiers from type <code>Tp</code>.  <a href="structcub_1_1_remove_qualifiers.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_sum.html">Sum</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Default sum functor.  <a href="structcub_1_1_sum.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_tex_iterator_r_a.html">TexIteratorRA</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple random-access iterator for loading primitive values through texture cache.  <a href="classcub_1_1_tex_iterator_r_a.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_tex_transform_iterator_r_a.html">TexTransformIteratorRA</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple random-access transform iterator for loading primitive values through texture cache and and subsequently applying a transformation operator.  <a href="classcub_1_1_tex_transform_iterator_r_a.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_traits.html">Traits</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">Type traits.  <a href="structcub_1_1_traits.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_transform_iterator_r_a.html">TransformIteratorRA</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A simple random-access transform iterator for applying a transformation operator.  <a href="classcub_1_1_transform_iterator_r_a.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">struct &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="structcub_1_1_uninitialized.html">Uninitialized</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">A storage-backing wrapper that allows types with non-trivial constructors to be aliased in unions.  <a href="structcub_1_1_uninitialized.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_reduce.html">WarpReduce</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_warp_reduce.html" title="The WarpReduce class provides collective methods for computing a parallel reduction of items partitio...">WarpReduce</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel reduction of items partitioned across CUDA warp threads. </p>
<div class="image">
<img src="warp_reduce_logo.png" alt="warp_reduce_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_warp_reduce.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:"><td class="memItemLeft" align="right" valign="top">class &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="classcub_1_1_warp_scan.html">WarpScan</a></td></tr>
<tr class="memdesc:"><td class="mdescLeft">&#160;</td><td class="mdescRight">The <a class="el" href="classcub_1_1_warp_scan.html" title="The WarpScan class provides collective methods for computing a parallel prefix scan of items partitio...">WarpScan</a> class provides <a href="index.html#sec0"><em>collective</em></a> methods for computing a parallel prefix scan of items partitioned across CUDA warp threads. </p>
<div class="image">
<img src="warp_scan_logo.png" alt="warp_scan_logo.png"/>
<div class="caption">
.</div></div>
  <a href="classcub_1_1_warp_scan.html#details">More...</a><br/></td></tr>
<tr class="separator:"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="enum-members"></a>
Enumerations</h2></td></tr>
<tr class="memitem:a845ca81ef9b346ef7429a35d4fdc80d2"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecub.html#a845ca81ef9b346ef7429a35d4fdc80d2">BlockHistogramAlgorithm</a> { <a class="el" href="namespacecub.html#a845ca81ef9b346ef7429a35d4fdc80d2ac14962247ae5693de2de53a8db40ca54">BLOCK_HISTO_SORT</a>, 
<a class="el" href="namespacecub.html#a845ca81ef9b346ef7429a35d4fdc80d2abfc89fce43fcda396d18a4b2801025aa">BLOCK_HISTO_ATOMIC</a>
 }</td></tr>
<tr class="memdesc:a845ca81ef9b346ef7429a35d4fdc80d2"><td class="mdescLeft">&#160;</td><td class="mdescRight">BlockHistogramAlgorithm enumerates alternative algorithms for the parallel construction of block-wide histograms.  <a href="namespacecub.html#a845ca81ef9b346ef7429a35d4fdc80d2">More...</a><br/></td></tr>
<tr class="separator:a845ca81ef9b346ef7429a35d4fdc80d2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a9d7e37497fdd99864c57adecda710401"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401">BlockLoadAlgorithm</a> { <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835">BLOCK_LOAD_DIRECT</a>, 
<a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a826be9d4df1c44c0e5c00a9c9c136965">BLOCK_LOAD_VECTORIZE</a>, 
<a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401acd94f285472e8f7c883a7407f6f4efc4">BLOCK_LOAD_TRANSPOSE</a>, 
<a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401aba21066c09e6c6ebcf08c7e86216eb1f">BLOCK_LOAD_WARP_TRANSPOSE</a>
 }</td></tr>
<tr class="memdesc:a9d7e37497fdd99864c57adecda710401"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401" title="cub::BlockLoadAlgorithm enumerates alternative algorithms for cub::BlockLoad to read a linear segment...">cub::BlockLoadAlgorithm</a> enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad</a> to read a linear segment of data from memory into a blocked arrangement across a CUDA thread block.  <a href="namespacecub.html#a9d7e37497fdd99864c57adecda710401">More...</a><br/></td></tr>
<tr class="separator:a9d7e37497fdd99864c57adecda710401"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:add0251c713859b8974806079e498d10a"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecub.html#add0251c713859b8974806079e498d10a">BlockReduceAlgorithm</a> { <a class="el" href="namespacecub.html#add0251c713859b8974806079e498d10aab32651e17a8a42207e74b7ed8d1aa4d2">BLOCK_REDUCE_RAKING</a>, 
<a class="el" href="namespacecub.html#add0251c713859b8974806079e498d10aa993903176f938273fa1ff5d4daa808e5">BLOCK_REDUCE_WARP_REDUCTIONS</a>
 }</td></tr>
<tr class="separator:add0251c713859b8974806079e498d10a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:abec44bba36037c547e7e84906d0d23ab"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab">BlockScanAlgorithm</a> { <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c">BLOCK_SCAN_RAKING</a>, 
<a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23abadc8dfe7760ac97ddde467ecf1070674f">BLOCK_SCAN_RAKING_MEMOIZE</a>, 
<a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba7f51e58246eb53f1a97bd1bc8c0f400f">BLOCK_SCAN_WARP_SCANS</a>
 }</td></tr>
<tr class="memdesc:abec44bba36037c547e7e84906d0d23ab"><td class="mdescLeft">&#160;</td><td class="mdescRight">BlockScanAlgorithm enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan</a> to compute a parallel prefix scan across a CUDA thread block.  <a href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab">More...</a><br/></td></tr>
<tr class="separator:abec44bba36037c547e7e84906d0d23ab"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:a839b145451e9eec3d44172e3c3619700"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700">BlockStoreAlgorithm</a> { <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880">BLOCK_STORE_DIRECT</a>, 
<a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a0ccd625a7f2f3649155cbd5a27adfb41">BLOCK_STORE_VECTORIZE</a>, 
<a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700ab0bbe20613466c3cedfcfea33a97d69c">BLOCK_STORE_TRANSPOSE</a>, 
<a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9178e3be2332f84a66761cee4ee6f3e3">BLOCK_STORE_WARP_TRANSPOSE</a>
 }</td></tr>
<tr class="memdesc:a839b145451e9eec3d44172e3c3619700"><td class="mdescLeft">&#160;</td><td class="mdescRight"><a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700" title="cub::BlockStoreAlgorithm enumerates alternative algorithms for cub::BlockStore to write a blocked arr...">cub::BlockStoreAlgorithm</a> enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore</a> to write a blocked arrangement of items across a CUDA thread block to a linear segment of memory.  <a href="namespacecub.html#a839b145451e9eec3d44172e3c3619700">More...</a><br/></td></tr>
<tr class="separator:a839b145451e9eec3d44172e3c3619700"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4733b6d40e923244502e6f5b200766ef"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___util_module.html#ga4733b6d40e923244502e6f5b200766ef">Category</a> { <b>NOT_A_NUMBER</b>, 
<b>SIGNED_INTEGER</b>, 
<b>UNSIGNED_INTEGER</b>, 
<b>FLOATING_POINT</b>
 }</td></tr>
<tr class="memdesc:ga4733b6d40e923244502e6f5b200766ef"><td class="mdescLeft">&#160;</td><td class="mdescRight">Basic type traits categories. <br/></td></tr>
<tr class="separator:ga4733b6d40e923244502e6f5b200766ef"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d">PtxLoadModifier</a> { <br/>
&#160;&#160;<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55daee682a9988eceaefd9a53a7ef92fe6e2">LOAD_DEFAULT</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55dacec87ab52dfb122d1555bfdfe44c23ae">LOAD_CA</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55dafad1bebebe5fe5ee6f6dfb5dc220009f">LOAD_CG</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da6f0ae5826ad5eee54046731f7c71c874">LOAD_CS</a>, 
<br/>
&#160;&#160;<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da83161be6f5901c6559583c0f7ea96076">LOAD_CV</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da0217758d7be24e5fd450d1d49aec36b8">LOAD_LDG</a>, 
<a class="el" href="group___io_module.html#gga023420f30fec7d4b187fc98f4fd2a55da390f7351eaf6b5b6ea23b4ec089a8236">LOAD_VOLATILE</a>
<br/>
 }</td></tr>
<tr class="memdesc:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Enumeration of PTX cache-modifiers for memory load operations.  <a href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d">More...</a><br/></td></tr>
<tr class="separator:ga023420f30fec7d4b187fc98f4fd2a55d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gae9c7d6a6af7104f528509182ac9c9da2"><td class="memItemLeft" align="right" valign="top">enum &#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2">PtxStoreModifier</a> { <br/>
&#160;&#160;<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a434aa4b3efc8e0c0ce4f1a00cdad26bd">STORE_DEFAULT</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2afd29f1e82f0a9dc8258fb7977b0a8237">STORE_WB</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2aacea07ea298b89dd1962a40b4823652d">STORE_CG</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a00ae8891d1acad179d134fdd60d7839b">STORE_CS</a>, 
<br/>
&#160;&#160;<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2a5bf33b361cb9f3229588c95402fc484f">STORE_WT</a>, 
<a class="el" href="group___io_module.html#ggae9c7d6a6af7104f528509182ac9c9da2aba537dcdd7a709bcdff169e014dcaf08">STORE_VOLATILE</a>
<br/>
 }</td></tr>
<tr class="memdesc:gae9c7d6a6af7104f528509182ac9c9da2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Enumeration of PTX cache-modifiers for memory store operations.  <a href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2">More...</a><br/></td></tr>
<tr class="separator:gae9c7d6a6af7104f528509182ac9c9da2"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table><table class="memberdecls">
<tr class="heading"><td colspan="2"><h2 class="groupheader"><a name="func-members"></a>
Functions</h2></td></tr>
<tr class="memitem:ga5a175d2a88f63f7f1ab30e8b4f2cfa95"><td class="memItemLeft" align="right" valign="top">__host__ __device__ <br class="typebreak"/>
__forceinline__ cudaError_t&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___util_module.html#ga5a175d2a88f63f7f1ab30e8b4f2cfa95">Debug</a> (cudaError_t error, const char *filename, int line)</td></tr>
<tr class="memdesc:ga5a175d2a88f63f7f1ab30e8b4f2cfa95"><td class="mdescLeft">&#160;</td><td class="mdescRight">CUB error reporting macro (prints error messages to stderr)  <a href="group___util_module.html#ga5a175d2a88f63f7f1ab30e8b4f2cfa95">More...</a><br/></td></tr>
<tr class="separator:ga5a175d2a88f63f7f1ab30e8b4f2cfa95"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gac3d86a53650d110862fe8a84fce61c2e"><td class="memItemLeft" align="right" valign="top"><a class="anchor" id="gac3d86a53650d110862fe8a84fce61c2e"></a>
__host__ __device__ <br class="typebreak"/>
__forceinline__ cudaError_t&#160;</td><td class="memItemRight" valign="bottom"><a class="el" href="group___util_module.html#gac3d86a53650d110862fe8a84fce61c2e">PtxVersion</a> (int &amp;ptx_version)</td></tr>
<tr class="memdesc:gac3d86a53650d110862fe8a84fce61c2e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Retrieves the PTX version (major * 100 + minor * 10) <br/></td></tr>
<tr class="separator:gac3d86a53650d110862fe8a84fce61c2e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Simple I/O</div></td></tr>
<tr class="memitem:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ <br class="typebreak"/>
std::iterator_traits<br class="typebreak"/>
&lt; InputIteratorRA &gt;<br class="typebreak"/>
::value_type&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga2467f2337aeeea3a3a1b39d19996d3c2">ThreadLoad</a> (InputIteratorRA itr)</td></tr>
<tr class="memdesc:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Thread utility for reading memory using <a class="el" href="group___io_module.html#ga023420f30fec7d4b187fc98f4fd2a55d" title="Enumeration of PTX cache-modifiers for memory load operations. ">cub::PtxLoadModifier</a> cache modifiers.  <a href="group___io_module.html#ga2467f2337aeeea3a3a1b39d19996d3c2">More...</a><br/></td></tr>
<tr class="separator:ga2467f2337aeeea3a3a1b39d19996d3c2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename OutputIteratorRA , typename T &gt; </td></tr>
<tr class="memitem:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga870b5b156d3d0a65a8c551884b18a2d6">ThreadStore</a> (OutputIteratorRA itr, T val)</td></tr>
<tr class="memdesc:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="mdescLeft">&#160;</td><td class="mdescRight">Thread utility for writing memory using <a class="el" href="group___io_module.html#gae9c7d6a6af7104f528509182ac9c9da2" title="Enumeration of PTX cache-modifiers for memory store operations. ">cub::PtxStoreModifier</a> cache modifiers.  <a href="group___io_module.html#ga870b5b156d3d0a65a8c551884b18a2d6">More...</a><br/></td></tr>
<tr class="separator:ga870b5b156d3d0a65a8c551884b18a2d6"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Sequential reduction over statically-sized array types</div></td></tr>
<tr class="memitem:ga48320d1f549b7931ae1f72040f9e42dc"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ReductionOp &gt; </td></tr>
<tr class="memitem:ga48320d1f549b7931ae1f72040f9e42dc"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga48320d1f549b7931ae1f72040f9e42dc">ThreadReduce</a> (T *input, ReductionOp reduction_op, T prefix)</td></tr>
<tr class="memdesc:ga48320d1f549b7931ae1f72040f9e42dc"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential reduction over <code>LENGTH</code> elements of the <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#ga48320d1f549b7931ae1f72040f9e42dc">More...</a><br/></td></tr>
<tr class="separator:ga48320d1f549b7931ae1f72040f9e42dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gaeafffc414146089dec3bf734a0022b23"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ReductionOp &gt; </td></tr>
<tr class="memitem:gaeafffc414146089dec3bf734a0022b23"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#gaeafffc414146089dec3bf734a0022b23">ThreadReduce</a> (T *input, ReductionOp reduction_op)</td></tr>
<tr class="memdesc:gaeafffc414146089dec3bf734a0022b23"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential reduction over <code>LENGTH</code> elements of the <code>input</code> array. The aggregate is returned.  <a href="group___thread_module.html#gaeafffc414146089dec3bf734a0022b23">More...</a><br/></td></tr>
<tr class="separator:gaeafffc414146089dec3bf734a0022b23"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gaa532e530f559bb9922754bd610f5fe84"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ReductionOp &gt; </td></tr>
<tr class="memitem:gaa532e530f559bb9922754bd610f5fe84"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#gaa532e530f559bb9922754bd610f5fe84">ThreadReduce</a> (T(&amp;input)[LENGTH], ReductionOp reduction_op, T prefix)</td></tr>
<tr class="memdesc:gaa532e530f559bb9922754bd610f5fe84"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential reduction over the statically-sized <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#gaa532e530f559bb9922754bd610f5fe84">More...</a><br/></td></tr>
<tr class="separator:gaa532e530f559bb9922754bd610f5fe84"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga5d4d3204c873db5af48a3a3f831a9ef8"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ReductionOp &gt; </td></tr>
<tr class="memitem:ga5d4d3204c873db5af48a3a3f831a9ef8"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga5d4d3204c873db5af48a3a3f831a9ef8">ThreadReduce</a> (T(&amp;input)[LENGTH], ReductionOp reduction_op)</td></tr>
<tr class="memdesc:ga5d4d3204c873db5af48a3a3f831a9ef8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Serial reduction with the specified operator.  <a href="group___thread_module.html#ga5d4d3204c873db5af48a3a3f831a9ef8">More...</a><br/></td></tr>
<tr class="separator:ga5d4d3204c873db5af48a3a3f831a9ef8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Sequential prefix scan over statically-sized array types</div></td></tr>
<tr class="memitem:ga37fa2b0ca9bb32096517480e98894545"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:ga37fa2b0ca9bb32096517480e98894545"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga37fa2b0ca9bb32096517480e98894545">ThreadScanExclusive</a> (T *input, T *output, ScanOp scan_op, T prefix, bool apply_prefix=true)</td></tr>
<tr class="memdesc:ga37fa2b0ca9bb32096517480e98894545"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential exclusive prefix scan over <code>LENGTH</code> elements of the <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#ga37fa2b0ca9bb32096517480e98894545">More...</a><br/></td></tr>
<tr class="separator:ga37fa2b0ca9bb32096517480e98894545"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga145ad5e3b05134aef8d86f84df9c25b8"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:ga145ad5e3b05134aef8d86f84df9c25b8"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga145ad5e3b05134aef8d86f84df9c25b8">ThreadScanExclusive</a> (T(&amp;input)[LENGTH], T(&amp;output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix=true)</td></tr>
<tr class="memdesc:ga145ad5e3b05134aef8d86f84df9c25b8"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential exclusive prefix scan over the statically-sized <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#ga145ad5e3b05134aef8d86f84df9c25b8">More...</a><br/></td></tr>
<tr class="separator:ga145ad5e3b05134aef8d86f84df9c25b8"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4004b2a6bfb2b4f024b0389e9c46074c"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:ga4004b2a6bfb2b4f024b0389e9c46074c"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga4004b2a6bfb2b4f024b0389e9c46074c">ThreadScanInclusive</a> (T *input, T *output, ScanOp scan_op)</td></tr>
<tr class="memdesc:ga4004b2a6bfb2b4f024b0389e9c46074c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential inclusive prefix scan over <code>LENGTH</code> elements of the <code>input</code> array. The aggregate is returned.  <a href="group___thread_module.html#ga4004b2a6bfb2b4f024b0389e9c46074c">More...</a><br/></td></tr>
<tr class="separator:ga4004b2a6bfb2b4f024b0389e9c46074c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gae78193279081a554432b6aea10a7295a"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:gae78193279081a554432b6aea10a7295a"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#gae78193279081a554432b6aea10a7295a">ThreadScanInclusive</a> (T(&amp;input)[LENGTH], T(&amp;output)[LENGTH], ScanOp scan_op)</td></tr>
<tr class="memdesc:gae78193279081a554432b6aea10a7295a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential inclusive prefix scan over the statically-sized <code>input</code> array. The aggregate is returned.  <a href="group___thread_module.html#gae78193279081a554432b6aea10a7295a">More...</a><br/></td></tr>
<tr class="separator:gae78193279081a554432b6aea10a7295a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga0e00b3a9b522249581b2ff98f670087b"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:ga0e00b3a9b522249581b2ff98f670087b"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#ga0e00b3a9b522249581b2ff98f670087b">ThreadScanInclusive</a> (T *input, T *output, ScanOp scan_op, T prefix, bool apply_prefix=true)</td></tr>
<tr class="memdesc:ga0e00b3a9b522249581b2ff98f670087b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential inclusive prefix scan over <code>LENGTH</code> elements of the <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#ga0e00b3a9b522249581b2ff98f670087b">More...</a><br/></td></tr>
<tr class="separator:ga0e00b3a9b522249581b2ff98f670087b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gadb4aed8a804827b6fb97eca790572099"><td class="memTemplParams" colspan="2">template&lt;int LENGTH, typename T , typename ScanOp &gt; </td></tr>
<tr class="memitem:gadb4aed8a804827b6fb97eca790572099"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ T&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___thread_module.html#gadb4aed8a804827b6fb97eca790572099">ThreadScanInclusive</a> (T(&amp;input)[LENGTH], T(&amp;output)[LENGTH], ScanOp scan_op, T prefix, bool apply_prefix=true)</td></tr>
<tr class="memdesc:gadb4aed8a804827b6fb97eca790572099"><td class="mdescLeft">&#160;</td><td class="mdescRight">Perform a sequential inclusive prefix scan over the statically-sized <code>input</code> array, seeded with the specified <code>prefix</code>. The aggregate is returned.  <a href="group___thread_module.html#gadb4aed8a804827b6fb97eca790572099">More...</a><br/></td></tr>
<tr class="separator:gadb4aed8a804827b6fb97eca790572099"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Blocked I/O</div></td></tr>
<tr class="memitem:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga0f245535ce8c023d6d85bcfc0bf9907e">LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier.  <a href="group___io_module.html#ga0f245535ce8c023d6d85bcfc0bf9907e">More...</a><br/></td></tr>
<tr class="separator:ga0f245535ce8c023d6d85bcfc0bf9907e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gae694d6fa9abb0ef67802d12d2c19516c">LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:gae694d6fa9abb0ef67802d12d2c19516c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="group___io_module.html#gae694d6fa9abb0ef67802d12d2c19516c">More...</a><br/></td></tr>
<tr class="separator:gae694d6fa9abb0ef67802d12d2c19516c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga4aa27e3c558ba1ba7c306c834ad67b0f">LoadBlocked</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements..  <a href="group___io_module.html#ga4aa27e3c558ba1ba7c306c834ad67b0f">More...</a><br/></td></tr>
<tr class="separator:ga4aa27e3c558ba1ba7c306c834ad67b0f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga0e69bdc5b96bad62c9501af85532fc7b">StoreBlocked</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier.  <a href="group___io_module.html#ga0e69bdc5b96bad62c9501af85532fc7b">More...</a><br/></td></tr>
<tr class="separator:ga0e69bdc5b96bad62c9501af85532fc7b"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gac84f07cf1f61e45b26290515e9b4d3ae">StoreBlocked</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="group___io_module.html#gac84f07cf1f61e45b26290515e9b4d3ae">More...</a><br/></td></tr>
<tr class="separator:gac84f07cf1f61e45b26290515e9b4d3ae"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Striped I/O</div></td></tr>
<tr class="memitem:gaa32b5c960e511b67901f4a5ba423048e"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gaa32b5c960e511b67901f4a5ba423048e"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gaa32b5c960e511b67901f4a5ba423048e">LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gaa32b5c960e511b67901f4a5ba423048e"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier.  <a href="group___io_module.html#gaa32b5c960e511b67901f4a5ba423048e">More...</a><br/></td></tr>
<tr class="separator:gaa32b5c960e511b67901f4a5ba423048e"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga5a9c12543d3c373375c0bf03e9385a1f">LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="group___io_module.html#ga5a9c12543d3c373375c0bf03e9385a1f">More...</a><br/></td></tr>
<tr class="separator:ga5a9c12543d3c373375c0bf03e9385a1f"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga020063c6e3b8ffb09a87a00fe6823381">LoadStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga020063c6e3b8ffb09a87a00fe6823381"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.  <a href="group___io_module.html#ga020063c6e3b8ffb09a87a00fe6823381">More...</a><br/></td></tr>
<tr class="separator:ga020063c6e3b8ffb09a87a00fe6823381"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gadce0282ac2ac679aa9d046e534da6af2"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:gadce0282ac2ac679aa9d046e534da6af2"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gadce0282ac2ac679aa9d046e534da6af2">StoreStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gadce0282ac2ac679aa9d046e534da6af2"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier.  <a href="group___io_module.html#gadce0282ac2ac679aa9d046e534da6af2">More...</a><br/></td></tr>
<tr class="separator:gadce0282ac2ac679aa9d046e534da6af2"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, int BLOCK_THREADS, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga5d9cc4924d2bd640be9cb86e4754ce8c">StoreStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="group___io_module.html#ga5d9cc4924d2bd640be9cb86e4754ce8c">More...</a><br/></td></tr>
<tr class="separator:ga5d9cc4924d2bd640be9cb86e4754ce8c"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Warp-striped I/O</div></td></tr>
<tr class="memitem:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gaea39c52d05d5f8d078ff76b154fad26d">LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gaea39c52d05d5f8d078ff76b154fad26d"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier.  <a href="group___io_module.html#gaea39c52d05d5f8d078ff76b154fad26d">More...</a><br/></td></tr>
<tr class="separator:gaea39c52d05d5f8d078ff76b154fad26d"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga4486da2d836785d42f48ae379e8dd85a"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga4486da2d836785d42f48ae379e8dd85a"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga4486da2d836785d42f48ae379e8dd85a">LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga4486da2d836785d42f48ae379e8dd85a"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range.  <a href="group___io_module.html#ga4486da2d836785d42f48ae379e8dd85a">More...</a><br/></td></tr>
<tr class="separator:ga4486da2d836785d42f48ae379e8dd85a"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga92a5618a110d7963bbc6f54637b904b1"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename InputIteratorRA &gt; </td></tr>
<tr class="memitem:ga92a5618a110d7963bbc6f54637b904b1"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga92a5618a110d7963bbc6f54637b904b1">LoadWarpStriped</a> (int linear_tid, InputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items, T oob_default)</td></tr>
<tr class="memdesc:ga92a5618a110d7963bbc6f54637b904b1"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a warp-striped arrangement across the thread block using the specified cache modifier, guarded by range, with a fall-back assignment of out-of-bound elements.  <a href="group___io_module.html#ga92a5618a110d7963bbc6f54637b904b1">More...</a><br/></td></tr>
<tr class="separator:ga92a5618a110d7963bbc6f54637b904b1"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga07c4d38e195128da47fb9a8c582072d7"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga07c4d38e195128da47fb9a8c582072d7"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga07c4d38e195128da47fb9a8c582072d7">StoreWarpStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:ga07c4d38e195128da47fb9a8c582072d7"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier.  <a href="group___io_module.html#ga07c4d38e195128da47fb9a8c582072d7">More...</a><br/></td></tr>
<tr class="separator:ga07c4d38e195128da47fb9a8c582072d7"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD, typename OutputIteratorRA &gt; </td></tr>
<tr class="memitem:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#ga7b3a3b0c99fbfb1afc76b5fde54ba3dc">StoreWarpStriped</a> (int linear_tid, OutputIteratorRA block_itr, T(&amp;items)[ITEMS_PER_THREAD], int valid_items)</td></tr>
<tr class="memdesc:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a warp-striped arrangement of data across the thread block into a linear segment of items using the specified cache modifier, guarded by range.  <a href="group___io_module.html#ga7b3a3b0c99fbfb1afc76b5fde54ba3dc">More...</a><br/></td></tr>
<tr class="separator:ga7b3a3b0c99fbfb1afc76b5fde54ba3dc"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr><td colspan="2"><div class="groupHeader">Blocked, vectorized I/O</div></td></tr>
<tr class="memitem:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memTemplParams" colspan="2">template&lt;PtxLoadModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gadeb7216e9de14b5bd22f952cd1775eda">LoadBlockedVectorized</a> (int linear_tid, T *block_ptr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gadeb7216e9de14b5bd22f952cd1775eda"><td class="mdescLeft">&#160;</td><td class="mdescRight">Load a linear segment of items into a blocked arrangement across the thread block using the specified cache modifier.  <a href="group___io_module.html#gadeb7216e9de14b5bd22f952cd1775eda">More...</a><br/></td></tr>
<tr class="separator:gadeb7216e9de14b5bd22f952cd1775eda"><td class="memSeparator" colspan="2">&#160;</td></tr>
<tr class="memitem:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memTemplParams" colspan="2">template&lt;PtxStoreModifier MODIFIER, typename T , int ITEMS_PER_THREAD&gt; </td></tr>
<tr class="memitem:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memTemplItemLeft" align="right" valign="top">__device__ __forceinline__ void&#160;</td><td class="memTemplItemRight" valign="bottom"><a class="el" href="group___io_module.html#gac805e3fdfaecc3be16f81b0af5c1e4db">StoreBlockedVectorized</a> (int linear_tid, T *block_ptr, T(&amp;items)[ITEMS_PER_THREAD])</td></tr>
<tr class="memdesc:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="mdescLeft">&#160;</td><td class="mdescRight">Store a blocked arrangement of items across a thread block into a linear segment of items using the specified cache modifier.  <a href="group___io_module.html#gac805e3fdfaecc3be16f81b0af5c1e4db">More...</a><br/></td></tr>
<tr class="separator:gac805e3fdfaecc3be16f81b0af5c1e4db"><td class="memSeparator" colspan="2">&#160;</td></tr>
</table>
<a name="details" id="details"></a><h2 class="groupheader">Detailed Description</h2>
<div class="textblock"><p>Optional outer namespace(s) </p>
<p>CUB namespace </p>
</div><h2 class="groupheader">Enumeration Type Documentation</h2>
<a class="anchor" id="a845ca81ef9b346ef7429a35d4fdc80d2"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="namespacecub.html#a845ca81ef9b346ef7429a35d4fdc80d2">cub::BlockHistogramAlgorithm</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>BlockHistogramAlgorithm enumerates alternative algorithms for the parallel construction of block-wide histograms. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="a845ca81ef9b346ef7429a35d4fdc80d2ac14962247ae5693de2de53a8db40ca54"></a>BLOCK_HISTO_SORT</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>Sorting followed by differentiation. Execution is comprised of two phases:<ol type="1">
<li>Sort the data using efficient radix sort</li>
<li>Look for "runs" of same-valued keys by detecting discontinuities; the run-lengths are histogram bin counts.</li>
</ol>
</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd>Delivers consistent throughput regardless of sample bin distribution. </dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a845ca81ef9b346ef7429a35d4fdc80d2abfc89fce43fcda396d18a4b2801025aa"></a>BLOCK_HISTO_ATOMIC</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>Use atomic addition to update byte counts directly</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd>Performance is strongly tied to the hardware implementation of atomic addition, and may be significantly degraded for non uniformly-random input distributions where many concurrent updates are likely to be made to the same bin counter. </dd></dl>
</td></tr>
</table>

<p>Definition at line <a class="el" href="block__histogram_8cuh_source.html#l00055">55</a> of file <a class="el" href="block__histogram_8cuh_source.html">block_histogram.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a9d7e37497fdd99864c57adecda710401"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401">cub::BlockLoadAlgorithm</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p><a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401" title="cub::BlockLoadAlgorithm enumerates alternative algorithms for cub::BlockLoad to read a linear segment...">cub::BlockLoadAlgorithm</a> enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_load.html" title="The BlockLoad class provides collective data movement methods for loading a linear segment of items f...">cub::BlockLoad</a> to read a linear segment of data from memory into a blocked arrangement across a CUDA thread block. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835"></a>BLOCK_LOAD_DIRECT</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of data is read directly from memory. The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub> reads the <em>i</em><sup>th</sup> segment of consecutive elements.</p>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) decreases as the access stride between threads increases (i.e., the number items per thread). </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a9d7e37497fdd99864c57adecda710401a826be9d4df1c44c0e5c00a9c9c136965"></a>BLOCK_LOAD_VECTORIZE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of data is read directly from memory using CUDA's built-in vectorized loads as a coalescing optimization. The thread block reads items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector loads to read the <em>i</em><sup>th</sup> segment of consecutive elements.</p>
<p>For example, <code>ld.global.v4.s32</code> instructions will be generated when <code>T</code> = <code>int</code> and <code>ITEMS_PER_THREAD</code> &gt; 4.</p>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high until the the access stride between threads (i.e., the number items per thread) exceeds the maximum vector load width (typically 4 items or 64B, whichever is lower).</li>
<li>The following conditions will prevent vectorization and loading will fall back to <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835">cub::BLOCK_LOAD_DIRECT</a>:<ul>
<li><code>ITEMS_PER_THREAD</code> is odd</li>
<li>The <code>InputIteratorRA</code> is not a simple pointer type</li>
<li>The block input offset is not quadword-aligned</li>
<li>The data type <code>T</code> is not a built-in primitive or CUDA vector type (e.g., <code>short</code>, <code>int2</code>, <code>double</code>, <code>float2</code>, etc.) </li>
</ul>
</li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a9d7e37497fdd99864c57adecda710401acd94f285472e8f7c883a7407f6f4efc4"></a>BLOCK_LOAD_TRANSPOSE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>striped arrangement</em></a> of data is read directly from memory and then is locally transposed into a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a>. The thread block reads items in a parallel "strip-mining" fashion: thread<sub><em>i</em></sub> reads items having stride <code>BLOCK_THREADS</code> between them. <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange</a> is then used to locally reorder the items into a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a>.</p>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high regardless of items loaded per thread.</li>
<li>The local reordering incurs slightly longer latencies and throughput than the direct <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835">cub::BLOCK_LOAD_DIRECT</a> and <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a826be9d4df1c44c0e5c00a9c9c136965">cub::BLOCK_LOAD_VECTORIZE</a> alternatives. </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a9d7e37497fdd99864c57adecda710401aba21066c09e6c6ebcf08c7e86216eb1f"></a>BLOCK_LOAD_WARP_TRANSPOSE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>warp-striped arrangement</em></a> of data is read directly from memory and then is locally transposed into a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a>. Each warp reads its own contiguous segment in a parallel "strip-mining" fashion: lane<sub><em>i</em></sub> reads items having stride <code>WARP_THREADS</code> between them. <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange</a> is then used to locally reorder the items into a <a href="index.html#sec5sec4"><em>blocked arrangement</em></a>.</p>
<dl class="section user"><dt>Usage Considerations</dt><dd><ul>
<li>BLOCK_THREADS must be a multiple of WARP_THREADS</li>
</ul>
</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high regardless of items loaded per thread.</li>
<li>The local reordering incurs slightly longer latencies and throughput than the direct <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a2d4d8900d7e697e9dac4062e97d3d835">cub::BLOCK_LOAD_DIRECT</a> and <a class="el" href="namespacecub.html#a9d7e37497fdd99864c57adecda710401a826be9d4df1c44c0e5c00a9c9c136965">cub::BLOCK_LOAD_VECTORIZE</a> alternatives. </li>
</ul>
</dd></dl>
</td></tr>
</table>

<p>Definition at line <a class="el" href="block__load_8cuh_source.html#l00474">474</a> of file <a class="el" href="block__load_8cuh_source.html">block_load.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="add0251c713859b8974806079e498d10a"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="namespacecub.html#add0251c713859b8974806079e498d10a">cub::BlockReduceAlgorithm</a></td>
        </tr>
      </table>
</div><div class="memdoc">
<p>BlockReduceAlgorithm enumerates alternative algorithms for parallel reduction across a CUDA threadblock. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="add0251c713859b8974806079e498d10aab32651e17a8a42207e74b7ed8d1aa4d2"></a>BLOCK_REDUCE_RAKING</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>An efficient "raking" reduction algorithm. Execution is comprised of three phases:<ol type="1">
<li>Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.</li>
<li>Upsweep sequential reduction in shared memory. Threads within a single warp rake across segments of shared partial reductions.</li>
<li>A warp-synchronous Kogge-Stone style reduction within the raking warp.</li>
</ol>
</dd></dl>
<dl class="section user"><dt></dt><dd><div class="image">
<img src="block_reduce.png" alt="block_reduce.png"/>
</div>
 <div class="centercaption"><code>BLOCK_REDUCE_RAKING</code> data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div></dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Although this variant may suffer longer turnaround latencies when the GPU is under-occupied, it can often provide higher overall throughput across the GPU when suitably occupied. </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="add0251c713859b8974806079e498d10aa993903176f938273fa1ff5d4daa808e5"></a>BLOCK_REDUCE_WARP_REDUCTIONS</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>A quick "tiled warp-reductions" reduction algorithm. Execution is comprised of four phases:<ol type="1">
<li>Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.</li>
<li>Compute a shallow, but inefficient warp-synchronous Kogge-Stone style reduction within each warp.</li>
<li>A propagation phase where the warp reduction outputs in each warp are updated with the aggregate from each preceding warp.</li>
</ol>
</dd></dl>
<dl class="section user"><dt></dt><dd><div class="image">
<img src="block_scan_warpscans.png" alt="block_scan_warpscans.png"/>
</div>
 <div class="centercaption"><code>BLOCK_REDUCE_WARP_REDUCTIONS</code> data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div></dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Although this variant may suffer lower overall throughput across the GPU because due to a heavy reliance on inefficient warp-reductions, it can often provide lower turnaround latencies when the GPU is under-occupied. </li>
</ul>
</dd></dl>
</td></tr>
</table>

<p>Definition at line <a class="el" href="block__reduce_8cuh_source.html#l00058">58</a> of file <a class="el" href="block__reduce_8cuh_source.html">block_reduce.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="abec44bba36037c547e7e84906d0d23ab"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23ab">cub::BlockScanAlgorithm</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p>BlockScanAlgorithm enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_scan.html" title="The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items par...">cub::BlockScan</a> to compute a parallel prefix scan across a CUDA thread block. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c"></a>BLOCK_SCAN_RAKING</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>An efficient "raking reduce-then-scan" prefix scan algorithm. Execution is comprised of five phases:<ol type="1">
<li>Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.</li>
<li>Upsweep sequential reduction in shared memory. Threads within a single warp rake across segments of shared partial reductions.</li>
<li>A warp-synchronous Kogge-Stone style exclusive scan within the raking warp.</li>
<li>Downsweep sequential exclusive scan in shared memory. Threads within a single warp rake across segments of shared partial reductions, seeded with the warp-scan output.</li>
<li>Downsweep sequential scan in registers (if threads contribute more than one input), seeded with the raking scan output.</li>
</ol>
</dd></dl>
<dl class="section user"><dt></dt><dd><div class="image">
<img src="block_scan_raking.png" alt="block_scan_raking.png"/>
</div>
 <div class="centercaption"><code>BLOCK_SCAN_RAKING</code> data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div></dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Although this variant may suffer longer turnaround latencies when the GPU is under-occupied, it can often provide higher overall throughput across the GPU when suitably occupied. </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="abec44bba36037c547e7e84906d0d23abadc8dfe7760ac97ddde467ecf1070674f"></a>BLOCK_SCAN_RAKING_MEMOIZE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>Similar to <a class="el" href="namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c">cub::BLOCK_SCAN_RAKING</a>, but with fewer shared memory reads at the expense of higher register pressure. Raking threads preserve their "upsweep" segment of values in registers while performing warp-synchronous scan, allowing the "downsweep" not to re-read them from shared memory. </dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="abec44bba36037c547e7e84906d0d23aba7f51e58246eb53f1a97bd1bc8c0f400f"></a>BLOCK_SCAN_WARP_SCANS</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>A quick "tiled warpscans" prefix scan algorithm. Execution is comprised of four phases:<ol type="1">
<li>Upsweep sequential reduction in registers (if threads contribute more than one input each). Each thread then places the partial reduction of its item(s) into shared memory.</li>
<li>Compute a shallow, but inefficient warp-synchronous Kogge-Stone style scan within each warp.</li>
<li>A propagation phase where the warp scan outputs in each warp are updated with the aggregate from each preceding warp.</li>
<li>Downsweep sequential scan in registers (if threads contribute more than one input), seeded with the raking scan output.</li>
</ol>
</dd></dl>
<dl class="section user"><dt></dt><dd><div class="image">
<img src="block_scan_warpscans.png" alt="block_scan_warpscans.png"/>
</div>
 <div class="centercaption"><code>BLOCK_SCAN_WARP_SCANS</code> data flow for a hypothetical 16-thread threadblock and 4-thread raking warp.</div></dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>Although this variant may suffer lower overall throughput across the GPU because due to a heavy reliance on inefficient warpscans, it can often provide lower turnaround latencies when the GPU is under-occupied. </li>
</ul>
</dd></dl>
</td></tr>
</table>

<p>Definition at line <a class="el" href="block__scan_8cuh_source.html#l00056">56</a> of file <a class="el" href="block__scan_8cuh_source.html">block_scan.cuh</a>.</p>

</div>
</div>
<a class="anchor" id="a839b145451e9eec3d44172e3c3619700"></a>
<div class="memitem">
<div class="memproto">
      <table class="memname">
        <tr>
          <td class="memname">enum <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700">cub::BlockStoreAlgorithm</a></td>
        </tr>
      </table>
</div><div class="memdoc">

<p><a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700" title="cub::BlockStoreAlgorithm enumerates alternative algorithms for cub::BlockStore to write a blocked arr...">cub::BlockStoreAlgorithm</a> enumerates alternative algorithms for <a class="el" href="classcub_1_1_block_store.html" title="The BlockStore class provides collective data movement methods for writing a blocked arrangement of i...">cub::BlockStore</a> to write a blocked arrangement of items across a CUDA thread block to a linear segment of memory. </p>
<table class="fieldtable">
<tr><th colspan="2">Enumerator</th></tr><tr><td class="fieldname"><em><a class="anchor" id="a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880"></a>BLOCK_STORE_DIRECT</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of data is written directly to memory. The thread block writes items in a parallel "raking" fashion: thread<sub><em>i</em></sub> writes the <em>i</em><sup>th</sup> segment of consecutive elements.</p>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) decreases as the access stride between threads increases (i.e., the number items per thread). </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a839b145451e9eec3d44172e3c3619700a0ccd625a7f2f3649155cbd5a27adfb41"></a>BLOCK_STORE_VECTORIZE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd></dd></dl>
<p>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> of data is written directly to memory using CUDA's built-in vectorized stores as a coalescing optimization. The thread block writes items in a parallel "raking" fashion: thread<sub><em>i</em></sub> uses vector stores to write the <em>i</em><sup>th</sup> segment of consecutive elements.</p>
<p>For example, <code>st.global.v4.s32</code> instructions will be generated when <code>T</code> = <code>int</code> and <code>ITEMS_PER_THREAD</code> &gt; 4.</p>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high until the the access stride between threads (i.e., the number items per thread) exceeds the maximum vector store width (typically 4 items or 64B, whichever is lower).</li>
<li>The following conditions will prevent vectorization and writing will fall back to <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880">cub::BLOCK_STORE_DIRECT</a>:<ul>
<li><code>ITEMS_PER_THREAD</code> is odd</li>
<li>The <code>OutputIteratorRA</code> is not a simple pointer type</li>
<li>The block output offset is not quadword-aligned</li>
<li>The data type <code>T</code> is not a built-in primitive or CUDA vector type (e.g., <code>short</code>, <code>int2</code>, <code>double</code>, <code>float2</code>, etc.) </li>
</ul>
</li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a839b145451e9eec3d44172e3c3619700ab0bbe20613466c3cedfcfea33a97d69c"></a>BLOCK_STORE_TRANSPOSE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> is locally transposed into a <a href="index.html#sec5sec4"><em>striped arrangement</em></a> which is then written to memory. More specifically, <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange</a> used to locally reorder the items into a <a href="index.html#sec5sec4"><em>striped arrangement</em></a>, after which the thread block writes items in a parallel "strip-mining" fashion: consecutive items owned by thread<sub><em>i</em></sub> are written to memory with stride <code>BLOCK_THREADS</code> between them.</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high regardless of items written per thread.</li>
<li>The local reordering incurs slightly longer latencies and throughput than the direct <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880">cub::BLOCK_STORE_DIRECT</a> and <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a0ccd625a7f2f3649155cbd5a27adfb41">cub::BLOCK_STORE_VECTORIZE</a> alternatives. </li>
</ul>
</dd></dl>
</td></tr>
<tr><td class="fieldname"><em><a class="anchor" id="a839b145451e9eec3d44172e3c3619700a9178e3be2332f84a66761cee4ee6f3e3"></a>BLOCK_STORE_WARP_TRANSPOSE</em>&nbsp;</td><td class="fielddoc">
<dl class="section user"><dt>Overview</dt><dd>A <a href="index.html#sec5sec4"><em>blocked arrangement</em></a> is locally transposed into a <a href="index.html#sec5sec4"><em>warp-striped arrangement</em></a> which is then written to memory. More specifically, <a class="el" href="classcub_1_1_block_exchange.html" title="The BlockExchange class provides collective methods for rearranging data partitioned across a CUDA th...">cub::BlockExchange</a> used to locally reorder the items into a <a href="index.html#sec5sec4"><em>warp-striped arrangement</em></a>, after which each warp writes its own contiguous segment in a parallel "strip-mining" fashion: consecutive items owned by lane<sub><em>i</em></sub> are written to memory with stride <code>WARP_THREADS</code> between them.</dd></dl>
<dl class="section user"><dt>Performance Considerations</dt><dd><ul>
<li>The utilization of memory transactions (coalescing) remains high regardless of items written per thread.</li>
<li>The local reordering incurs slightly longer latencies and throughput than the direct <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a9b8dcc7b6b06bcfc24af4f499523b880">cub::BLOCK_STORE_DIRECT</a> and <a class="el" href="namespacecub.html#a839b145451e9eec3d44172e3c3619700a0ccd625a7f2f3649155cbd5a27adfb41">cub::BLOCK_STORE_VECTORIZE</a> alternatives. </li>
</ul>
</dd></dl>
</td></tr>
</table>

<p>Definition at line <a class="el" href="block__store_8cuh_source.html#l00364">364</a> of file <a class="el" href="block__store_8cuh_source.html">block_store.cuh</a>.</p>

</div>
</div>
</div><!-- contents -->
<!-- HTML footer for doxygen 1.8.3.1-->
<!-- start footer part -->
<hr class="footer"/><address class="footer"><small>
Generated on Fri Aug 23 2013 17:31:11 for CUB by &#160;<a href="http://www.doxygen.org/index.html">
<img class="footer" src="doxygen.png" alt="doxygen"/>
</a> 1.8.4
<br>
&copy; 2013 NVIDIA Corporation
</small></address>
</body>
</html>
